﻿//https://www.cnblogs.com/dwdxdy/p/3528711.html

#include <stdio.h>
#include<stdlib.h>
#include <opencv2/opencv.hpp>
//下面两个头文件基本都会用到，别省了
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

/*
二、单独使用Cuda API编程

利用Cuda Runtime API、Cuda Driver API实现一些操作的并行加速，使用过程需要管理CPU与GPU之间的数据传输，内核函数调用参数的设置，内核函数的优化等。

优点是处理过程受控于用户，用户可以实现更多的并行加速处理操作。

缺点是使用复杂，代码编写量较多，需要熟悉Cuda相关资料和API接口。下面是简单的示例程序：
*/
/*
__global__ void swap_rb_kernel(const uchar3* src,uchar3* dst,int width,int height)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.x + blockIdx.y * blockDim.y;
    
    if(x < width && y < height)
    {
        uchar3 v = src[y * width + x];
        dst[y * width + x].x = v.z;
        dst[y * width + x].y = v.y;
        dst[y * width + x].z = v.x;
    }
}

void swap_rb_caller(const uchar3* src,uchar3* dst,int width,int height)
{
    dim3 block(32,8);
    dim3 grid((width + block.x - 1)/block.x,(height + block.y - 1)/block.y);
    
    swap_rb_kernel<<<grid,block,0>>>(src,dst,width,height);
    cudaThreadSynchronize();
}

int main()
{
    Mat image = imread("lena.jpg");
    imshow("src",image);
    
    size_t memSize = image.cols*image.rows*sizeof(uchar3);
    uchar3* d_src = NULL;
    uchar3* d_dst = NULL;
    CUDA_SAFE_CALL(cudaMalloc((void**)&d_src,memSize));
    CUDA_SAFE_CALL(cudaMalloc((void**)&d_dst,memSize));
    CUDA_SAFE_CALL(cudaMempcy(d_src,image.data,memSize,cudaMemcpyHostToDevice));
    
    swap_rb_caller(d_src,d_dst,image.cols,image.rows);
    
    CUDA_SAFE_CALL(cudaMempcy(image.data,d_dst,memSize,cudaMemcpyDeviceToHost));
    imshow("gpu",image);
    waitKey(0);
    
    CUDA_SAFE_CALL(cudaFree(d_src));
    CUDA_SAFE_CALL(cudaFree(d_dst));
    return 0;
}
*/
/*
三、利用OpenCV中提供接口，并结合Cuda API编程

　　利用OpenCV已经提供的部分接口，完成一些Cuda编程的基本处理，简化编程的复杂程度；只是根据自己业务需求，自定义内核函数或扩展OpenCV已提供的内核函数。这样既可以充分利用OpenCV的特性，又可以满足业务的不同需求，使用方便，且易于扩展。下面是简单的示例程序：
  */


